#include "cuImage_filters.h"

#include "cuError.h"

texture<uchar4, 2, cudaReadModeNormalizedFloat> g_SrcTex;


__global__ 
void Kernel_Gaussian(uchar4 * pDst, float sigma, uint32 w, uint32 h)
{
  uint32 tidx = threadIdx.x + blockIdx.x * blockDim.x;
  uint32 tidy = threadIdx.y + blockIdx.y * blockDim.y;

  if (tidx < w && tidy < h)
  {
    float r = 2.0f * sigma + 1.0f;

    float4 cSum = {0.0f, 0.0f, 0.0f, 1.0f};
    float  wSum = 0.0f;


    for (int ix = -r; ix <= r; ix++)
      for (int iy = -r; iy <= r; iy++)
      {
        float4 cxy = tex2D(g_SrcTex, 
                   tidx + ix,
                   tidy + iy);

        float w = exp( -(ix*ix + iy*iy) / (sigma * sigma) );

        cSum.x += cxy.x * w;
        cSum.y += cxy.y * w;
        cSum.z += cxy.z * w;

        wSum += w;
      }

      cSum.x /= wSum;
      cSum.y /= wSum;
      cSum.z /= wSum;

      pDst[tidx + tidy * w] = make_uchar4((uchar) cSum.x * 255, 
                        (uchar) cSum.y * 255, 
                        (uchar) cSum.z * 255, 
                        (uchar) 255 );

  }
}

__global__ 
void Kernel_Gaussian_Tex2D(uchar4 * pDst, float sigma, uint32 w, uint32 h)
{
}


bool Wrapper_Gaussian(cuImage_RGBA8 *pDst, 
                      cuImage_RGBA8 *pSrc, 
                      float          sigma, 
                      uint8          rFlags)
{
  cudaError_t cuda_error;

  // if input pointers are NULL OR
  if (pDst == NULL || pSrc == NULL  )
    return false;

  // if output buffer is NULL
  if (pDst->getDPtr() == NULL)        
    return false;

  // if input memory is NULL AND if input cuda array is also NULL
  if (pSrc->getDPtr() == NULL && pSrc->getAPtr() == NULL)          
    return false;

  // if input width is different from output OR height is different
  if (pSrc->getP() != pDst->getP() || pSrc->getH() != pDst->getH() )

  if (rFlags & CU_IMAGE_A_MEMORY)
  {
    if (pSrc->getAPtr() == NULL)
      return false;

    CUDA_CHECK_ERROR( cudaBindTextureToArray(g_SrcTex, pSrc->getAPtr(), cudaCreateChannelDesc<uchar4>()) );

    uint32 w = pDst->getP();
    uint32 h = pDst->getH();

    uint32 threadsTotal = 128;
    uint32 warpsize = cuImage_RGBA8::device_properties.warpSize;
    dim3 threads(warpsize, threadsTotal / warpsize);
    dim3 blocks( (w + threads.x - 1) / threads.x,
           (h + threads.y - 1) / threads.y );

    Kernel_Gaussian<<<blocks, threads>>>(pDst->getDPtr(), sigma, w, h);

    CUDA_CHECK_ERROR( cudaThreadSynchronize() );

    return cuda_error == cudaSuccess;
  }

  if (rFlags & CU_IMAGE_D_MEMORY)
  {
    if (pSrc->getDPtr() == NULL)
      return false;

    CUDA_CHECK_ERROR( cudaBindTexture2D(NULL, g_SrcTex, 
                      pSrc->getDPtr(), 
                      cudaCreateChannelDesc<uchar4>(), 
                      pSrc->getW(), pSrc->getH(), 
                      pSrc->getP() * sizeof(uchar4)) );

    uint32 w = pDst->getP();
    uint32 h = pDst->getH();

    uint32 threadsTotal = 128;
    uint32 warpsize = cuImage_RGBA8::device_properties.warpSize;
    dim3 threads(warpsize, threadsTotal / warpsize);
    dim3 blocks( (w + threads.x - 1) / threads.x,  
           (h + threads.y - 1) / threads.y );

    Kernel_Gaussian<<<blocks, threads>>>(pDst->getDPtr(), sigma, w, h);

    CUDA_CHECK_ERROR( cudaThreadSynchronize() );

    return cuda_error == cudaSuccess;
  }

  return false;
}